<!DOCTYPE html>
<!--[if IE 8]><html class="no-js lt-ie9" lang="en" > <![endif]-->
<!--[if gt IE 8]><!--> <html class="no-js" lang="en" > <!--<![endif]-->
<head>
  <meta charset="utf-8">
  <meta http-equiv="X-UA-Compatible" content="IE=edge">
  <meta name="viewport" content="width=device-width, initial-scale=1.0">
  
  <meta name="author" content="PG-Strom Development Team">
  <link rel="shortcut icon" href="../img/favicon.ico">
  <title>PL/CUDA - PG-Strom Manual</title>
  <link href='https://fonts.googleapis.com/css?family=Lato:400,700|Roboto+Slab:400,700|Inconsolata:400,700' rel='stylesheet' type='text/css'>

  <link rel="stylesheet" href="../css/theme.css" type="text/css" />
  <link rel="stylesheet" href="../css/theme_extra.css" type="text/css" />
  <link rel="stylesheet" href="//cdnjs.cloudflare.com/ajax/libs/highlight.js/9.12.0/styles/github.min.css">
  <link href="//fonts.googleapis.com/earlyaccess/notosansjp.css" rel="stylesheet">
  <link href="//fonts.googleapis.com/css?family=Open+Sans:600,800" rel="stylesheet">
  <link href="../custom.css" rel="stylesheet">
  
  <script>
    // Current page data
    var mkdocs_page_name = "PL/CUDA";
    var mkdocs_page_input_path = "plcuda.md";
    var mkdocs_page_url = null;
  </script>
  
  <script src="../js/jquery-2.1.1.min.js" defer></script>
  <script src="../js/modernizr-2.8.3.min.js" defer></script>
  <script src="//cdnjs.cloudflare.com/ajax/libs/highlight.js/9.12.0/highlight.min.js"></script>
  <script>hljs.initHighlightingOnLoad();</script> 
  
</head>

<body class="wy-body-for-nav" role="document">

  <div class="wy-grid-for-nav">

    
    <nav data-toggle="wy-nav-shift" class="wy-nav-side stickynav">
      <div class="wy-side-nav-search">
        <a href=".." class="icon icon-home"> PG-Strom Manual</a>
        <div role="search">
  <form id ="rtd-search-form" class="wy-form" action="../search.html" method="get">
    <input type="text" name="q" placeholder="Search docs" />
  </form>
  [<strong>Japanese</strong> | <a href="../../plcuda/"    style="color: #cccccc">English</a>]
</div>
      </div>

      <div class="wy-menu wy-menu-vertical" data-spy="affix" role="navigation" aria-label="main navigation">
	<ul class="current">
	  
          
            <li class="toctree-l1">
		
    <a class="" href="..">はじめに</a>
	    </li>
          
            <li class="toctree-l1">
		
    <a class="" href="../install/">インストール</a>
	    </li>
          
            <li class="toctree-l1">
		
    <span class="caption-text">利用ガイド</span>
    <ul class="subnav">
                <li class="">
                    
    <a class="" href="../operations/">基本的な操作</a>
                </li>
                <li class="">
                    
    <a class="" href="../sys_admin/">システム管理</a>
                </li>
                <li class="">
                    
    <a class="" href="../brin/">インデックス対応</a>
                </li>
                <li class="">
                    
    <a class="" href="../partition/">パーティション</a>
                </li>
                <li class="">
                    
    <a class="" href="../troubles/">トラブルシューティング</a>
                </li>
    </ul>
	    </li>
          
            <li class="toctree-l1">
		
    <span class="caption-text">先進機能</span>
    <ul class="subnav">
                <li class="">
                    
    <a class="" href="../ssd2gpu/">SSDtoGPUダイレクトSQL</a>
                </li>
                <li class="">
                    
    <a class="" href="../arrow_fdw/">Arrow_fdw</a>
                </li>
                <li class="">
                    
    <a class="" href="../gstore_fdw/">Gstore_fdw</a>
                </li>
                <li class=" current">
                    
    <a class="current" href="./">PL/CUDA</a>
    <ul class="subnav">
            
    <li class="toctree-l3"><a href="#plcuda">PL/CUDA概要</a></li>
    

    <li class="toctree-l3"><a href="#plcuda_1">PL/CUDAの構造</a></li>
    

    <li class="toctree-l3"><a href="#plcuda_2">PL/CUDAリファレンス</a></li>
    
        <ul>
        
            <li><a class="toctree-l4" href="#plcuda_3">PL/CUDAの得意不得意</a></li>
        
            <li><a class="toctree-l4" href="#plcuda_4">PL/CUDAディレクティブ</a></li>
        
            <li><a class="toctree-l4" href="#plcuda_5">PL/CUDA 関連関数</a></li>
        
        </ul>
    

    </ul>
                </li>
    </ul>
	    </li>
          
            <li class="toctree-l1">
		
    <span class="caption-text">リファレンス</span>
    <ul class="subnav">
                <li class="">
                    
    <a class="" href="../ref_types/">データ型</a>
                </li>
                <li class="">
                    
    <a class="" href="../ref_devfuncs/">関数と演算子</a>
                </li>
                <li class="">
                    
    <a class="" href="../ref_sqlfuncs/">SQLオブジェクト</a>
                </li>
                <li class="">
                    
    <a class="" href="../ref_params/">GUCパラメータ</a>
                </li>
    </ul>
	    </li>
          
            <li class="toctree-l1">
		
    <a class="" href="../release_note/">リリースノート</a>
	    </li>
          
        </ul>
      </div>
      &nbsp;
    </nav>

    <section data-toggle="wy-nav-shift" class="wy-nav-content-wrap">

      
      <nav class="wy-nav-top" role="navigation" aria-label="top navigation">
        <i data-toggle="wy-nav-top" class="fa fa-bars"></i>
        <a href="..">PG-Strom Manual</a>
      </nav>

      
      <div class="wy-nav-content">
        <div class="rst-content">
          <div role="navigation" aria-label="breadcrumbs navigation">
  <ul class="wy-breadcrumbs">
    <li><a href="..">Docs</a> &raquo;</li>
    
      
        
          <li>先進機能 &raquo;</li>
        
      
    
    <li>PL/CUDA</li>
    <li class="wy-breadcrumbs-aside">
      
    </li>
  </ul>
  <hr/>
</div>
          <div role="main">
            <div class="section">
              
                <p>本章では、PL/CUDA言語を用いて、GPUで実行可能なネイティブプログラムをSQL関数として実装する方法について説明します。</p>
<h1 id="plcuda">PL/CUDA概要</h1>
<p>内部的に、PG-StromはSQL構文を元にCUDA言語によるGPUプログラムを生成し、これを実行時コンパイルによってGPU用命令バイナリを生成します。 CUDAとはNVIDIA社の提供するプログラミング環境で、C言語に似た構文を用いてGPUで並列実行可能なプログラムを記述する事ができます。 SQL構文からCUDAプログラムへの変換プロセスは内部的なもので、ユーザの視点からは、どのようなGPU用プラグラムが生成、実行されるのかを意識する必要はありません。</p>
<p>一方、PostgreSQLでは<code>CREATE LANGUAGE</code>構文を用いてSQL関数の記述に用いるプログラミング言語を追加する事ができます。 PL/CUDAとは<code>CREATE LANGUAGE</code>構文に対応した言語ハンドラで、SQLを元にPG-Stromが自動生成するGPUプログラムだけでなく、ユーザが実装した任意のGPUプログラムをSQL関数として実行する事が可能となります。</p>
<p>SQL関数の引数には、数値型や文字列型、行列型など、PG-Stromのサポートするデータ型を使用する事ができますが、これらはPL/CUDA実行系が自動的にGPU側へデータを転送するため、データベースとGPU間のデータロードについて意識する必要はありません。また同様に、PL/CUDA関数の戻り値（可変長データ型である場合を含む）もGPU側からCPU側へと書き戻され、SQL関数の戻り値として整形されます。</p>
<p>また、PL/CUDA関数の引数としてgstore_fdwを用いて定義した外部表を使用する事ができます。この場合、データは既にGPUにロード済みであるためPL/CUDA関数呼び出しのたびにデータロードを行う必要はなく、またPostgreSQL可変長データの長さ制限である1GBよりも大きなデータを使用する事ができます。</p>
<p>これらの特徴により、ユーザはGPUやデータベースとの間のデータの入出力といった定型的な処理に煩わされる事なく、統計解析ロジックの実装や高速化といった生産的な作業に注力する事ができます。</p>
<p><img alt="PL/CUDA Overview" src="../img/plcuda-overview.png" /></p>
<p><code>CREATE FUNCTION</code>構文を用いてPL/CUDA関数を定義すると、この関数の定義部を含むCUDAプログラムのソースコードを作成し、これをターゲットGPU向けにビルドします。
このCUDAプログラムは、引数の受け渡しと結果を返却するための補助的なコードを含む以外は、一般的なCUDAランタイムを用いたソフトウェアと全く同一で、CUDAの提供する各種のライブラリをインクルード／リンクする事も可能です。</p>
<p>PL/CUDA関数を用いて作成したネイティブのCUDAプログラムは、PostgreSQLバックエンドの子プロセスとして実行されます。
したがって、PostgreSQLとは独立したアドレス空間と、OSやGPUのリソースを持つ事になります。
CUDAプログラムには、ホストシステム上で実行されるホストコードと、GPU上で実行されるデバイスコードを含みます。ホストコードはC言語でプログラミング可能なあらゆるロジックを実行可能ですので、セキュリティ上の観点から、PL/CUDA関数の定義はデータベース特権ユーザに限定されています。</p>
<p>以下に単純なPL/CUDA関数の例を示します。 この関数は、同じ長さの<code>read</code>型配列を二つ引数に取り、そのドット積を<code>float</code>型で返却します。</p>
<pre><code>CREATE OR REPLACE FUNCTION
gpu_dot_product(real[], real[])
RETURNS float
AS $$
#plcuda_decl
#include &quot;cuda_matrix.h&quot;

KERNEL_FUNCTION_MAXTHREADS(void)
gpu_dot_product(double *p_dot,
                VectorTypeFloat *X,
                VectorTypeFloat *Y)
{
    size_t      index = get_global_id();
    size_t      nitems = X-&gt;height;
    float       v[MAXTHREADS_PER_BLOCK];
    float       sum;

    if (index &lt; nitems)
        v[get_local_id()] = X-&gt;values[index] * Y-&gt;values[index];
    else
        v[get_local_id()] = 0.0;

    sum = pgstromTotalSum(v, MAXTHREADS_PER_BLOCK);
    if (get_local_id() == 0)
        atomicAdd(p_dot, (double)sum);
    __syncthreads();
}
#plcuda_begin
{
    size_t      nitems;
    int         blockSz;
    int         gridSz;
    double     *dot;
    cudaError_t rc;

    if (!VALIDATE_ARRAY_VECTOR_TYPE_STRICT(arg1, PG_FLOAT4OID) ||
        !VALIDATE_ARRAY_VECTOR_TYPE_STRICT(arg2, PG_FLOAT4OID))
        EEXIT(&quot;arguments are not vector like array&quot;);
    nitems = ARRAY_VECTOR_HEIGHT(arg1);
    if (nitems != ARRAY_VECTOR_HEIGHT(arg2))
        EEXIT(&quot;length of arguments mismatch&quot;);

    rc = cudaMallocManaged(&amp;dot, sizeof(double));
    if (rc != cudaSuccess)
        CUEXIT(rc, &quot;failed on cudaMallocManaged&quot;);
    memset(dot, 0, sizeof(double));

    blockSz = MAXTHREADS_PER_BLOCK;
    gridSz = (nitems + MAXTHREADS_PER_BLOCK - 1) / MAXTHREADS_PER_BLOCK;
    gpu_dot_product&lt;&lt;&lt;gridSz,blockSz&gt;&gt;&gt;(dot,
                                        (VectorTypeFloat *)arg1,
                                        (VectorTypeFloat *)arg2);
    rc = cudaStreamSynchronize(NULL);
    if (rc != cudaSuccess)
        CUEXIT(rc, &quot;failed on cudaStreamSynchronize&quot;);

    return *dot;
}
#plcuda_end
$$ LANGUAGE 'plcuda';
</code></pre>

<p>PL/CUDA実行系は、<code>#plcuda_begin</code>と<code>#plcuda_end</code>で囲まれた部分にSQL関数の引数の受け渡しを行う処理を付加して、CUDAプログラムのエントリポイントを作成します。
<code>#plcuda_decl</code>と<code>#plcuda_begin</code>で囲まれた部分は、GPUデバイス関数やその他のホスト関数を宣言するためのブロックで、ソースコード上では上記のエントリポイントより前に配置されます。</p>
<p>CUDAプログラムのエントリポイントでは、<code>arg1</code>、<code>arg2</code>、... という形でSQL関数の引数を参照する事ができます。</p>
<p>上記の例では、<code>real[]</code>配列型である<code>arg1</code>および<code>arg2</code>がエントリポイントへ渡され、<code>VALIDATE_ARRAY_VECTOR_TYPE_STRICT</code>マクロによってNULLを含まない32bit浮動小数点型の１次元配列であるかどうかを検証しています。</p>
<p>返り値も同様に、SQLデータ型に相当するCUDA Cデータ型をエントリポイントから返します。
エントリポイントが<code>return</code>で値を返さない場合（または明示的に<code>exit()</code>で終了コード1を返した場合）、PL/CUDA関数は<code>NULL</code>値を返却したものとして扱われます。</p>
<p>上記のサンプルプログラムでは、SQL関数から受け取った<code>real</code>型配列を検証した後、<code>cudaMallocManaged</code>で結果バッファを獲得した後、GPUカーネル関数である<code>gpu_dot_product</code>を呼出してドット積を計算しています。</p>
<p>この関数の実行結果は以下の通りです。ランダムに生成した10,000個の要素を持つベクトル同士の内積を計算しています。</p>
<pre><code>postgres=# SELECT gpu_dot_product(array_matrix(random()::real),
                                  array_matrix(random()::real))
             FROM generate_series(1,10000);
 gpu_dot_product
------------------
 3.71461999509484
(1 row)
</code></pre>

<h1 id="plcuda_1">PL/CUDAの構造</h1>
<p>PL/CUDAの関数定義は、<code>#plcuda_decl</code>、<code>#plcuda_begin</code>、および<code>#plcuda_end</code>の各ディレクティブによって分割されたコードブロックから構成されます。各コードブロックには各々の目的に応じたユーザ定義のCUDA Cコードを記述する事ができ、これらは、PL/CUDA言語ハンドラとの引数及び結果の受け渡しを行うロジックと結合し一個のソースファイルへと再構成されます。</p>
<pre><code>#plcuda_decl
  [...any declarations...]
#plcuda_begin
  [...host code in the entrypoint...]
#plcuda_end
</code></pre>

<p><code>#plcuda_decl</code>より始まるコードブロックは、<code>__host__</code>および<code>__device__</code>属性を持つCUDA C関数や変数を完全な形で記述する事ができます。
このコードブロックは、最終的に構築されるソースファイル上で、<code>#plcuda_begin</code>...<code>#plcuda_end</code>ブロックを含むエントリポイントよりも前方に位置します。</p>
<p>また、CUDA Cの<code>#include</code>構文を用いて外部ヘッダファイルをインクルードする場合には、このコードブロックに記述するようにして下さい。</p>
<p><code>#plcuda_begin</code>より始まるコードブロックは、ホストコードであるエントリポイント関数の一部として組み込まれます。したがって、関数名や引数の型などを記述する事はできません。
エントリポイント関数は、当該コードブロックに制御が移る前に、パイプを介してSQL関数の引数をPostgreSQLバックエンドから受信し、<code>arg1</code>、<code>arg2</code>、…という名前で参照できるようセットアップを行います。</p>
<p>これらの変数は、SQLデータ型に応じて以下の表に示すCUDA Cとしての表現を持ちます。</p>
<table>
<thead>
<tr>
<th align="center">SQLデータ型</th>
<th align="center">CUDA Cデータ型</th>
<th align="center">説明</th>
</tr>
</thead>
<tbody>
<tr>
<td align="center"><code>reggstore</code></td>
<td align="center"><code>void *</code></td>
<td align="center">Gstore_fdw外部表のOID</td>
</tr>
<tr>
<td align="center"><code>real</code></td>
<td align="center"><code>float</code></td>
<td align="center">32bit浮動小数点型</td>
</tr>
<tr>
<td align="center"><code>float</code></td>
<td align="center"><code>double</code></td>
<td align="center">64bit浮動小数点型</td>
</tr>
<tr>
<td align="center">その他のインライン型</td>
<td align="center"><code>Datum</code></td>
<td align="center"><code>int</code>、<code>date</code>など</td>
</tr>
<tr>
<td align="center">固定長ポインタ型</td>
<td align="center"><code>void *</code></td>
<td align="center"><code>uuid</code>など</td>
</tr>
<tr>
<td align="center">可変長データ型(varlena)</td>
<td align="center"><code>varlena *</code></td>
<td align="center"><code>text</code>、<code>real[]</code>など</td>
</tr>
</tbody>
</table>
<p>PL/CUDA言語ハンドラは、上記のコードブロックから一個のCUDA Cソースファイルを作成し、宣言時または実行時に１度だけ<code>nvcc</code>コンパイラでこれをビルドしCUDAプログラムを生成します。
<code>#plcuda_include</code>ディレクティブを含む場合はCUDA Cソースファイルが実行時にしか確定しないため、実行時にのみこれをビルドします。ただし、同一内容のCUDAプログラムがビルド済みである場合にはこれを再利用します。</p>
<p><img alt="PL/CUDA Callflow" src="../img/plcuda-callflow.png" /></p>
<p>SQLからPL/CUDA関数を呼び出すと、PL/CUDA言語ハンドラはビルド済みのCUDAプログラムを起動し、パイプを通じてSQL関数の引数をコピーします。引数はCUDAプログラム内の引数バッファに格納され、これらは<code>arg1</code>や<code>arg2</code>などの名前で参照する事が可能です。</p>
<p>可変長データ型などCUDA Cプログラム上でポインタとして表現されるデータ型は、引数バッファへの参照として初期化されます。引数バッファは<code>cudaMallocManaged()</code>によって獲得されたmanaged memory領域であるため、当該ポインタはホスト⇔デバイス間の明示的なDMAなしに使用する事ができます。</p>
<p>引数が<code>reggstore</code>型を持つ場合は特殊です。これは本来Gstore_Fdw外部テーブルのOID（4バイト整数）を表現するデータ型ですが、PL/CUDAの引数として与えられた場合はGstore_fdwが獲得しているGPUデバイスメモリへの参照へと置き換えられます。
引数は<code>GstoreIpcMapping</code>オブジェクトへの参照として初期化され、<code>GstoreIpcMapping::map</code>にはGstore_Fdw外部テーブルの確保したGPUデバイスメモリをマップしたアドレスが入ります。
当該領域を物理的に保持しているGPUデバイスIDは<code>GstoreIpcHandle::device_id</code>を、当該領域の長さは<code>GstoreIpcHandle::rawsize</code>を参照してください。</p>
<pre><code>typedef struct
{
    cl_uint     __vl_len;       /* 4B varlena header */
    cl_short    device_id;      /* GPU device where pinning on */
    cl_char     format;         /* one of GSTORE_FDW_FORMAT__* */
    cl_char     __padding__;    /* reserved */
    cl_long     rawsize;        /* length in bytes */
    union {
#ifdef CU_IPC_HANDLE_SIZE
        CUipcMemHandle      d;  /* CUDA driver API */
#endif
#ifdef CUDA_IPC_HANDLE_SIZE
        cudaIpcMemHandle_t  r;  /* CUDA runtime API */
#endif
        char                data[64];
    } ipc_mhandle;
} GstoreIpcHandle;

typedef struct
{
    GstoreIpcHandle h; /* IPChandle of Gstore_Fdw */
    void       *map;    /* mapped device pointer */
} GstoreIpcMapping;
</code></pre>

<p>PL/CUDA関数の処理結果を返すには、SQLデータ型に対応するCUDA Cのデータをエントリポイントから<code>return</code>で返却します。
値を明示的に返却しない場合、CUDA Cでのデータ型がポインタ型であり<code>NULL</code>を返した場合、あるいはCUDAプログラムが<code>exit(1)</code>によりステータスコード1で終了した場合は、PL/CUDA関数はSQLに対して<code>null</code>を返したものとして扱われます。</p>
<h1 id="plcuda_2">PL/CUDAリファレンス</h1>
<p>本節はPL/CUDA関数のディレクティブ、および関連するSQL関数のリファレンスです。</p>
<h2 id="plcuda_3">PL/CUDAの得意不得意</h2>
<p>PL/CUDA関数が呼び出されると、その背後でCUDAプログラムが起動され、CUDAプログラムはGPUデバイスの初期化を行います。これらの一連の処理は決して軽いものではなく、例えば単純なスカラー値の比較を行うようなロジックをPL/CUDA関数で実装し、10億行のフルテーブルと同時に使用するという使い方は推奨されません。</p>
<p>一方で、ひとたびGPUデバイスの初期化が完了すれば、GPUの持つ数千プロセッサコアを利用して大量データを高速に処理する事が可能です。特に、繰り返し計算により最適パラメータを計算する機械学習や統計解析のように、ワークロードに占める計算の割合が大きな問題に適すると言えるでしょう。</p>
<p>処理すべきデータが増加すると、CUDAプログラムとのデータの受け渡し方法にも注意が必要です。
PostgreSQLは配列型をサポートしており、整数型や実数型のデータを高々数百万個程度受け渡すのであれば手軽な方法です。
しかし、配列型を含むPostgreSQLの可変長データの最大長は1GBであるため、これより巨大なデータの受け渡しにはデータの分割など工夫が必要です。また、SQL関数の引数をセットアップするのはPostgreSQLバックエンドプロセスで、この処理はシングルスレッドで動作するためGB単位のメモリ操作には相応の時間を要します。</p>
<p>データサイズが数百MBを越えてきた段階で、Gstore_Fdw外部テーブルの利用を検討してください。
Gstore_Fdwを通して予めGPUデバイスメモリにデータをロードする事で、PL/CUDA関数の呼び出し時に長大な引数をセットアップする必要がなく、また、GPUデバイスメモリ容量が許す限り、GBを越えるサイズのデータを保持する事が可能です。</p>
<h2 id="plcuda_4">PL/CUDAディレクティブ</h2>
<h3 id="plcuda_decl"><code>#plcuda_decl</code></h3>
<p>このディレクティブは<code>__host__</code>および<code>__device__</code>属性を持つCUDA C関数や変数を含むコードブロックを開始します。PL/CUDA言語ハンドラは、CUDAプログラムのソースファイル上で、このコードブロックをエントリポイントよりも前にそのままコピーします。</p>
<p>このディレクティブの使用は任意ですが、エントリポイントから呼び出すべきGPUカーネル関数を宣言しなければPL/CUDA関数を使用する意味はありませんので、通常は一つ以上のGPUカーネル関数を含む事になります。</p>
<h3 id="plcuda_begin"><code>#plcuda_begin</code></h3>
<p>このディレクティブは、CUDAプログラムのエントリポイントを構成するコードブロックを開始します。
CUDAプログラムは、受け取ったPL/CUDA関数の引数を<code>arg1</code>、<code>arg2</code>、...という変数名で参照可能となるよう初期化を行った上で、コードブロックへと制御を移します。当該コードブロックはホストコードであり、CPUで動作する制御ロジックや、GPUカーネルを呼出しての計算処理を記述する事ができます。</p>
<p>結果を返すには、CUDA Cの<code>return</code>構文でPL/CUDA関数の返り値に応じたデータを返します。</p>
<h3 id="plcuda_end"><code>#plcuda_end</code></h3>
<p>コードブロックの終了を宣言します。 なお、あるコードブロックの内側で他のコードブロックの開始を宣言した場合、現在のコードブロックは暗黙のうちに<code>#plcuda_end</code>ディレクティブによって終了したものとして扱われます。</p>
<h3 id="plcuda_include-ltfunction-namegt"><code>#plcuda_include &lt;function name&gt;</code></h3>
<p>このディレクティブはCUDA Cの<code>#include</code>と似ていますが、ヘッダファイルではなく、指定されたSQL関数の実行結果をディレクティブの存在していた場所に挿入します。
オプションで指定するSQL関数はPL/CUDA関数と同一の引数をとり<code>text</code>型を返す必要があります。</p>
<p>これは例えば、大量データ間の類似度を計算する際に、計算のアルゴリズムはほとんど同一であるにも関わらず距離計算のロジックだけが異なるバリエーションを動的に作り出す事が可能で、PL/CUDA関数の保守を簡素化する事ができます。</p>
<h3 id="plcuda_library-ltlibrary-namegt"><code>#plcuda_library &lt;library name&gt;</code></h3>
<p>CUDAプログラムをビルドする際にリンクするライブラリ名を指定します。
<code>&lt;library name&gt;</code>に記述するのは、<code>nvcc</code>コマンドの<code>-l</code>オプションに相当する文字列です。
例えば<code>libcublas.so</code>ライブラリをリンクする場合には、接頭語の<code>lib</code>と拡張子の<code>.so</code>を省略した<code>cublas</code>と指定します。
現在のところ、CUDA Toolkitの標準ライブラリパス（<code>/usr/local/cuda/lib64</code>）にインストールされたライブラリのみを指定する事ができます。</p>
<h3 id="plcuda_sanity_check-ltfunctiongt"><code>#plcuda_sanity_check &lt;function&gt;</code></h3>
<p>GPUカーネルの起動に先立って、引数の妥当性を検証するためのSQL関数をしています。
デフォルトでは妥当性検証関数は設定されていません。
GPUデバイスの初期化などを行う必要があるため、通常、GPUカーネル関数の起動はCPU上で別の関数を起動するよりも重い処理です。もし引数がPL/CUDA関数の仕様からは許容できない値を持っている場合、GPUカーネル関数を実行する数千～数百万（場合によってはそれ以上の）のGPUスレッドは、ただ引数の妥当性をチェックしてエラー状態を返却するためだけに起動されます。GPUカーネル関数を実行する前に、引数の妥当性チェックを十分小さなコストで行えるならば、妥当性検証関数を使用してGPUカーネル関数の実行前にエラーを発生させることを考慮すべきです。 妥当性検証関数は、PL/CUDA関数と同じ型の引数を持ち、<code>bool</code>返す関数です。</p>
<h2 id="plcuda_5">PL/CUDA 関連関数</h2>
<table>
<thead>
<tr>
<th align="left">関数定義</th>
<th align="center">結果型</th>
<th align="left">説明</th>
</tr>
</thead>
<tbody>
<tr>
<td align="left"><code>plcuda_function_source(regproc)</code></td>
<td align="center"><code>text</code></td>
<td align="left">引数としてPL/CUDA関数のOIDを与えると、PL/CUDA関数から生成されるGPUカーネルのソースコードを返します。</td>
</tr>
</tbody>
</table>
<h3 id="plcuda_6">PL/CUDA関数呼び出し支援</h3>
<p>以下の関数群は、PL/CUDA関数の呼び出しを簡便にするために提供されています。</p>
<table>
<thead>
<tr>
<th align="left">関数定義</th>
<th align="center">結果型</th>
<th align="left">説明</th>
</tr>
</thead>
<tbody>
<tr>
<td align="left"><code>attnums_of(regclass,text[])</code></td>
<td align="center"><code>smallint[]</code></td>
<td align="left">第一引数で指定したテーブルの第二引数で指定した列名（複数可）の列番号を配列として返します。</td>
</tr>
<tr>
<td align="left"><code>attnum_of(regclass,text)</code></td>
<td align="center"><code>smallint</code></td>
<td align="left">第一引数で指定したテーブルの第二引数で指定した列名の列番号を返します。</td>
</tr>
<tr>
<td align="left"><code>atttypes_of(regclass,text[])</code></td>
<td align="center"><code>regtype[]</code></td>
<td align="left">第一引数で指定したテーブルの第二引数で指定した列名（複数可）のデータ型を配列として返します。</td>
</tr>
<tr>
<td align="left"><code>atttype_of(regclass,text)</code></td>
<td align="center"><code>regtype</code></td>
<td align="left">第一引数で指定したテーブルの第二引数で指定した列名のデータ型を返します。</td>
</tr>
<tr>
<td align="left"><code>attrs_types_check(regclass,text[],regtype[])</code></td>
<td align="center"><code>bool</code></td>
<td align="left">第一引数で指定したテーブルの、第二引数で指定した列名（複数可）のデータ型が、第三引数で指定したデータ型とそれぞれ一致しているかどうかを調べます。</td>
</tr>
<tr>
<td align="left"><code>attrs_type_check(regclass,text[],regtype)</code></td>
<td align="center"><code>bool</code></td>
<td align="left">第一引数で指定したテーブルの、第二引数で指定した列名（複数可）のデータ型が、全て第三引数で指定したデータ型と一致しているかどうかを調べます。</td>
</tr>
</tbody>
</table>
<h3 id="_1">配列ベースの行列型関数</h3>
<p>本節ではPG-Stromの提供する配列ベースの行列型をサポートするSQL関数について説明します。</p>
<p>PostgreSQLには行列を表現するための専用のデータ型は存在していませんが、以下の条件を満たす二次元配列をあたかも行列であるかのように取り扱う事が可能です。</p>
<ul>
<li>二次元配列である</li>
<li>各次元の配列要素が1から始まる</li>
<li>NULL値を含まない</li>
<li>配列の大きさが1GBを越えない。（PostgreSQL可変長データ表現による制約）</li>
<li><code>smallint</code>、<code>int</code>、<code>bigint</code>、<code>real</code>または<code>float</code>型の配列である</li>
</ul>
<p>配列がこれらの条件を満たす時、行列の(i,j)要素の位置は添え字から一意に特定する事ができ、GPUスレッドが自らの処理すべきデータを効率的に取り出す事を可能とします。また、通常の行形式データとは異なり、計算に必要なデータのみをロードする事になるため、メモリ消費やデータ転送の点で有利です。 PG-Stromは、この様な疑似的な行列型をサポートするため、以下に示すSQL関数を提供しています。</p>
<table>
<thead>
<tr>
<th align="left">関数定義</th>
<th align="center">結果型</th>
<th align="left">説明</th>
</tr>
</thead>
<tbody>
<tr>
<td align="left"><code>array_matrix(variadic arg, ...)</code></td>
<td align="center"><code>array</code></td>
<td align="left">入力された行を全て連結した配列ベース行列を返す集約関数です。例えば、<code>float</code>型の引数x、y、zを1000行入力すると、同じ<code>float</code>型で3列×1000行の配列ベース行列を返します。<br>この関数は可変長引数を取るよう定義されており、<code>arg</code>は1個以上の<code>smallint</code>、<code>int</code>、<code>bigint</code>、<code>real</code>または<code>float</code>型のスカラー値で、全ての<code>arg</code>値は同じデータ型を持つ必要があります。</td>
</tr>
<tr>
<td align="left"><code>matrix_unnest(array)</code></td>
<td align="center"><code>record</code></td>
<td align="left">配列ベース行列を行の集合に展開する集合関数です。<code>array</code>は<code>smallint</code>、<code>int</code>、<code>bigint</code>、<code>real</code>または<code>float</code>型の配列で、行列の幅に応じて1個以上のカラムからなる<code>record</code>型を返却します。例えば、10列×500行から成る行列の場合、各レコードは行列要素のデータ型を持つ10個のカラムからなり、これが500行生成されます。 <br>標準の<code>unnest</code>関数と似ていますが、<code>record</code>型を生成するため、<code>AS (colname1 type[, ...])</code>句を用いて返却されるべきレコードの型を指定する必要があります。</td>
</tr>
<tr>
<td align="left"><code>rbind(array, array)</code></td>
<td align="center"><code>array</code></td>
<td align="left"><code>array</code>は<code>smallint</code>、<code>int</code>、<code>bigint</code>、<code>real</code>または<code>float</code>型の配列です。<br>二つの配列ベース行列を縦方向に結合します。双方の行列は同一の要素データ型を持つ必要があり、行列の幅が等しくない場合は足りない部分を0で埋めます。</td>
</tr>
<tr>
<td align="left"><code>rbind(array)</code></td>
<td align="center"><code>array</code></td>
<td align="left"><code>array</code>は<code>smallint</code>、<code>int</code>、<code>bigint</code>、<code>real</code>または<code>float</code>型の配列です。<code>rbind(array, array)</code>と似ていますが、集合関数として動作し入力された全ての配列ベース行列を縦方向に結合します。</td>
</tr>
<tr>
<td align="left"><code>cbind(array, array)</code></td>
<td align="center"><code>array</code></td>
<td align="left"><code>array</code>は<code>smallint</code>、<code>int</code>、<code>bigint</code>、<code>real</code>または<code>float</code>型の配列で、二つの配列ベース行列を横方向に結合します。双方の行列は同一の要素データ型を持つ必要があり、行列の高さ等しくない場合は足りない部分を0で埋めます。</td>
</tr>
<tr>
<td align="left"><code>cbind(array)</code></td>
<td align="center"><code>array</code></td>
<td align="left"><code>array</code>は<code>smallint</code>、<code>int</code>、<code>bigint</code>、<code>real</code>または<code>float</code>型の配列で、<code>cbind(array, array)</code>と似ていますが、集合関数として動作し入力された全ての配列ベース行列を横方向に結合します。</td>
</tr>
<tr>
<td align="left"><code>transpose(array)</code></td>
<td align="center"><code>array</code></td>
<td align="left"><code>array</code>は<code>smallint</code>、<code>int</code>、<code>bigint</code>、<code>real</code>または<code>float</code>型の配列で、行列の幅と高さが入れ替わった転置行列を生成します。</td>
</tr>
<tr>
<td align="left"><code>array_matrix_validation(anyarray)</code></td>
<td align="center"><code>bool</code></td>
<td align="left">入力された配列（<code>anyarray</code>）が、配列ベース行列として妥当かどうかを検査します。 PL/CUDA関数実行前の引数の妥当性検証や、DOMAIN型を定義する時の検査制約としての利用を想定しています。</td>
</tr>
<tr>
<td align="left"><code>array_matrix_height(array)</code></td>
<td align="center"><code>int</code></td>
<td align="left"><code>array</code>は<code>smallint</code>、<code>int</code>、<code>bigint</code>、<code>real</code>または<code>float</code>型の配列で、配列ベース行列の高さを返却します。</td>
</tr>
<tr>
<td align="left"><code>array_matrix_width(array)</code></td>
<td align="center"><code>int</code></td>
<td align="left"><code>array</code>は<code>smallint</code>、<code>int</code>、<code>bigint</code>、<code>real</code>または<code>float</code>型の配列で、配列ベース行列の幅を返却します。</td>
</tr>
</tbody>
</table>
              
            </div>
          </div>
          <footer>
  
    <div class="rst-footer-buttons" role="navigation" aria-label="footer navigation">
      
        <a href="../ref_types/" class="btn btn-neutral float-right" title="データ型">Next <span class="icon icon-circle-arrow-right"></span></a>
      
      
        <a href="../gstore_fdw/" class="btn btn-neutral" title="Gstore_fdw"><span class="icon icon-circle-arrow-left"></span> Previous</a>
      
    </div>
  

  <hr/>

  <div role="contentinfo">
    <!-- Copyright etc -->
    
  </div>

  Built with <a href="http://www.mkdocs.org">MkDocs</a> using a <a href="https://github.com/snide/sphinx_rtd_theme">theme</a> provided by <a href="https://readthedocs.org">Read the Docs</a>.
</footer>
      
        </div>
      </div>

    </section>

  </div>

  <div class="rst-versions" role="note" style="cursor: pointer">
    <span class="rst-current-version" data-toggle="rst-current-version">
      
      
        <span><a href="../gstore_fdw/" style="color: #fcfcfc;">&laquo; Previous</a></span>
      
      
        <span style="margin-left: 15px"><a href="../ref_types/" style="color: #fcfcfc">Next &raquo;</a></span>
      
    </span>
</div>
    <script>var base_url = '..';</script>
    <script src="../js/theme.js" defer></script>
      <script src="../search/main.js" defer></script>

</body>
</html>
